#pragma once 
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

#include "../cvtColor.h"

#define RGB2Y(R, G, B)  ( 16  + 0.183f * (R) + 0.614f * (G) + 0.062f * (B) )
#define RGB2U(R, G, B)  ( 128 - 0.101f * (R) - 0.339f * (G) + 0.439f * (B) )
#define RGB2V(R, G, B)  ( 128 + 0.439f * (R) - 0.399f * (G) - 0.040f * (B) )

#define YUV2R(Y, U, V) ( 1.164f *((Y) - 16) + 1.792f * ((V) - 128) )
#define YUV2G(Y, U, V) ( 1.164f *((Y) - 16) - 0.213f *((U) - 128) - 0.534f *((V) - 128) )
#define YUV2B(Y, U, V) ( 1.164f *((Y) - 16) + 2.114f *((U) - 128))

#define CLIPVALUE(x, minValue, maxValue) ((x) < (minValue) ? (minValue) : ((x) > (maxValue) ? (maxValue) : (x)))

namespace Image2D
{

__global__ static void __rgbToNv12 (
    const unsigned char* dpRgb, size_t rgbPitch,
    unsigned char*       dpY,   size_t yPitch,
    unsigned char*       dpUv,  size_t uvPitch,
    int width,                  int height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    unsigned char r = dpRgb[h * rgbPitch + w * 3 + 0];
    unsigned char g = dpRgb[h * rgbPitch + w * 3 + 1];
    unsigned char b = dpRgb[h * rgbPitch + w * 3 + 2];

    dpY[h * yPitch + w] = (unsigned char)(CLIPVALUE(RGB2Y(r, g, b), 0, 255));
    if (h % 2 == 0 && w % 2 == 0)
    {
        dpUv[(h / 2) * uvPitch + w / 2 * 2 + 0] = (unsigned char)(CLIPVALUE(RGB2U(r, g, b), 0, 255));
        dpUv[(h / 2) * uvPitch + w / 2 * 2 + 1] = (unsigned char)(CLIPVALUE(RGB2V(r, g, b), 0, 255));
    }
}

__global__ static void __rgbToYuv420p(
    const unsigned char* dpRgb, size_t rgbPitch,
    unsigned char*       dpY,   size_t yPitch,
    unsigned char*       dpU,   size_t uPitch,
    unsigned char*       dpV,   size_t vPitch,
    int                  width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    unsigned char r = dpRgb[h * rgbPitch + w * 3 + 0];
    unsigned char g = dpRgb[h * rgbPitch + w * 3 + 1];
    unsigned char b = dpRgb[h * rgbPitch + w * 3 + 2];

    dpY[h * yPitch + w] = (unsigned char)(CLIPVALUE(RGB2Y(r, g, b), 0, 255));
    if (h % 2 == 0 && w % 2 == 0)
    {
        dpU[(h / 2) * uPitch + w / 2] = (unsigned char)(CLIPVALUE(RGB2U(r, g, b), 0, 255));
        dpV[(h / 2) * vPitch + w / 2] = (unsigned char)(CLIPVALUE(RGB2V(r, g, b), 0, 255));
    }
}

__global__ static void __rgbToBgr    (
    const unsigned char* dpRgb, size_t rgbPitch,
    unsigned char*       dpBgr, size_t bgrPitch,
    int                  width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    dpBgr[h * bgrPitch + w * 3 + 0] = dpRgb[h * rgbPitch + w * 3 + 2];
    dpBgr[h * bgrPitch + w * 3 + 1] = dpRgb[h * rgbPitch + w * 3 + 1];
    dpBgr[h * bgrPitch + w * 3 + 2] = dpRgb[h * rgbPitch + w * 3 + 0];
}

__global__ static void __nv12ToRgb    (
    const unsigned char*  dpY,   size_t yPitch,
    const unsigned char*  dpUv,  size_t uvPitch,
    unsigned char*        dpRgb, size_t rgbPitch,
    int                   width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    unsigned char y = dpY [ h *      yPitch +  w];
    unsigned char u = dpUv[(h / 2) * uvPitch + w / 2 * 2 + 0];
    unsigned char v = dpUv[(h / 2) * uvPitch + w / 2 * 2 + 1];

    dpRgb[h * rgbPitch + w * 3 + 0] = (unsigned char)CLIPVALUE(YUV2R(y, u, v), 0, 255);
    dpRgb[h * rgbPitch + w * 3 + 1] = (unsigned char)CLIPVALUE(YUV2G(y, u, v), 0, 255);
    dpRgb[h * rgbPitch + w * 3 + 2] = (unsigned char)CLIPVALUE(YUV2B(y, u, v), 0, 255);
}

__global__ static void __nv12ToYuv420p(
    const unsigned char*  dpY,    size_t yPitch,
    const unsigned char*  dpUv,   size_t uvPitch,
    unsigned char*        dpY4,   size_t y4Pitch,
    unsigned char*        dpU4,   size_t u4Pitch,
    unsigned char*        dpV4,   size_t v4Pitch,
    int                   width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    dpY4[ h *      y4Pitch + w]      = dpY[  h *      yPitch +  w];
    dpU4[(h / 2) * u4Pitch + w / 2 ] = dpUv[(h / 2) * uvPitch + w / 2  * 2 + 0];
    dpV4[(h / 2) * v4Pitch + w / 2 ] = dpUv[(h / 2) * uvPitch + w / 2  * 2 + 1];
}

__global__ static void __nv12ToBgr    (
    const unsigned char*  dpY,   size_t yPitch,
    const unsigned char*  dpUv,  size_t uvPitch,
    unsigned char*        dpBgr, size_t bgrPitch,
    int                   width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    unsigned char y = dpY [ h *      yPitch +  w];
    unsigned char u = dpUv[(h / 2) * uvPitch + w / 2 * 2 + 0];
    unsigned char v = dpUv[(h / 2) * uvPitch + w / 2 * 2 + 1];

    dpBgr[h * bgrPitch + w * 3 + 2] = (unsigned char)CLIPVALUE(YUV2R(y, u, v), 0, 255);
	dpBgr[h * bgrPitch + w * 3 + 1] = (unsigned char)CLIPVALUE(YUV2G(y, u, v), 0, 255);
	dpBgr[h * bgrPitch + w * 3 + 0] = (unsigned char)CLIPVALUE(YUV2B(y, u, v), 0, 255);
}

__global__ static void __yuv420pToNv12(
    const unsigned char* dpY4,   size_t y4Pitch,
    const unsigned char* dpU4,   size_t u4Pitch,
    const unsigned char* dpV4,   size_t v4Pitch,
    unsigned char*       dpY,    size_t yPitch,
    unsigned char*       dpUv,   size_t uvPitch,
    int                   width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    dpY [ h *      yPitch  + w]              = dpY4[ h *      y4Pitch + w];
    dpUv[(h / 2) * uvPitch + w / 2  * 2 + 0] = dpU4[(h / 2) * u4Pitch + w / 2 ];
    dpUv[(h / 2) * uvPitch + w / 2  * 2 + 1] = dpV4[(h / 2) * v4Pitch + w / 2 ];
}

__global__ static void __yuv420pToRgb(
    const unsigned char* dpY,   size_t yPitch,
    const unsigned char* dpU,   size_t uPitch,
    const unsigned char* dpV,   size_t vPitch,
    unsigned char*       dpRgb, size_t rgbPitch,
    int                  width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    unsigned char y = dpY[ h *      yPitch + w];
    unsigned char u = dpU[(h / 2) * uPitch + w / 2];
    unsigned char v = dpV[(h / 2) * vPitch + w / 2];

    dpRgb[h * rgbPitch + w * 3 + 0] = (unsigned char)CLIPVALUE(YUV2R(y, u, v), 0, 255);
    dpRgb[h * rgbPitch + w * 3 + 1] = (unsigned char)CLIPVALUE(YUV2G(y, u, v), 0, 255);
    dpRgb[h * rgbPitch + w * 3 + 2] = (unsigned char)CLIPVALUE(YUV2B(y, u, v), 0, 255);
}

__global__ static void __yuv420pToBgr (
    const unsigned char* dpY,   size_t yPitch,
    const unsigned char* dpU,   size_t uPitch,
    const unsigned char* dpV,   size_t vPitch,
    unsigned char*       dpBgr, size_t bgrPitch,
    int                  width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    unsigned char y = dpY[ h *      yPitch + w];
    unsigned char u = dpU[(h / 2) * uPitch + w / 2];
    unsigned char v = dpV[(h / 2) * vPitch + w / 2];

    dpBgr[h * bgrPitch + w * 3 + 2] = (unsigned char)CLIPVALUE(YUV2R(y, u, v), 0, 255);;
    dpBgr[h * bgrPitch + w * 3 + 1] = (unsigned char)CLIPVALUE(YUV2G(y, u, v), 0, 255);;
    dpBgr[h * bgrPitch + w * 3 + 0] = (unsigned char)CLIPVALUE(YUV2B(y, u, v), 0, 255);;
}

__global__ static void __bgrToNv12   (
    const unsigned char* dpBgr, size_t bgrPitch,
    unsigned char*       dpY,   size_t yPitch,
    unsigned char*       dpUv,  size_t uvPitch,
    int width,                  int height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    unsigned char r = dpBgr[h * bgrPitch + w * 3 + 2];
    unsigned char g = dpBgr[h * bgrPitch + w * 3 + 1];
    unsigned char b = dpBgr[h * bgrPitch + w * 3 + 0];

    dpY[h * yPitch + w] = (unsigned char)(CLIPVALUE(RGB2Y(r, g, b), 0, 255));
    if (h % 2 == 0 && w % 2 == 0)
    {
        dpUv[(h / 2) * uvPitch + w / 2 * 2 + 0] = (unsigned char)(CLIPVALUE(RGB2U(r, g, b), 0, 255));
        dpUv[(h / 2) * uvPitch + w / 2 * 2 + 1] = (unsigned char)(CLIPVALUE(RGB2V(r, g, b), 0, 255));
    }
}

__global__ static void __bgrToYuv420p(
    const unsigned char* dpBgr, size_t bgrPitch,
    unsigned char*       dpY,   size_t yPitch,
    unsigned char*       dpU,   size_t uPitch,
    unsigned char*       dpV,   size_t vPitch,
    int                  width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    unsigned char r = dpBgr[h * bgrPitch + w * 3 + 2];
    unsigned char g = dpBgr[h * bgrPitch + w * 3 + 1];
    unsigned char b = dpBgr[h * bgrPitch + w * 3 + 0];

    dpY[h * yPitch + w] = (unsigned char)(CLIPVALUE(RGB2Y(r, g, b), 0, 255));
    if (h % 2 == 0 && w % 2 == 0)
    {
        dpU[(h / 2) * uPitch + w / 2] = (unsigned char)(CLIPVALUE(RGB2U(r, g, b), 0, 255));
        dpV[(h / 2) * vPitch + w / 2] = (unsigned char)(CLIPVALUE(RGB2V(r, g, b), 0, 255));
    }
}

__global__ static void __bgrToRgb (
    const unsigned char*  dpBgr, size_t bgrPitch,
    unsigned char*        dpRgb, size_t rgbPitch,
    int                   width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    dpRgb[h * rgbPitch + w * 3 + 2] = dpBgr[h * bgrPitch + w * 3 + 0];
    dpRgb[h * rgbPitch + w * 3 + 1] = dpBgr[h * bgrPitch + w * 3 + 1];
    dpRgb[h * rgbPitch + w * 3 + 0] = dpBgr[h * bgrPitch + w * 3 + 2];
}

__global__ static void __rgbaToBgra(
    const unsigned char*  dpRgba, size_t rgbaPitch,
    unsigned char*        dpBgra, size_t bgraPitch,
    int                   width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    dpBgra[h * bgraPitch + w * 4 + 0] = dpRgba[w * rgbaPitch + w * 4 + 2];
    dpBgra[h * bgraPitch + w * 4 + 1] = dpRgba[w * rgbaPitch + w * 4 + 1];
    dpBgra[h * bgraPitch + w * 4 + 2] = dpRgba[w * rgbaPitch + w * 4 + 0];
    dpBgra[h * bgraPitch + w * 4 + 3] = dpRgba[w * rgbaPitch + w * 4 + 3];
}

__global__ static void __bgraToRgba(
    const unsigned char* dpBgra, size_t bgraPitch,
    unsigned char*       dpRgba, size_t rgbaPitch,
    int                  width, int    height
)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    if (w >= width || h >= height)
        return;

    dpRgba[h * rgbaPitch + w * 4 + 2] = dpBgra[h * bgraPitch + w * 4 + 0];
    dpRgba[h * rgbaPitch + w * 4 + 1] = dpBgra[h * bgraPitch + w * 4 + 1];
    dpRgba[h * rgbaPitch + w * 4 + 0] = dpBgra[h * bgraPitch + w * 4 + 2];
    dpRgba[h * rgbaPitch + w * 4 + 4] = dpBgra[h * bgraPitch + w * 4 + 4];
}

}